#ifndef CL_TKIT_H
#define CL_TKIT_H

#include "Log.h"

#include <iostream>
#include <sstream>
#include <list>
#include <vector>
#include "cldeviceinfo.h"
#include <string.h>
#include "clcheck.h"

using namespace std;

class CLTkit
{
public:
    static cl_device_id deviceid(cl_uint platformNum, cl_uint deviceNum);
    static cl_program load_program(cl_context context, const char *filename, cl_uint count = 1);
    static std::vector<cl_platform_id> platformids();
    static std::vector<cl_device_id> deviceids(const  cl_platform_id &);
    static void getNDRangeWorkItemSize(cl_uint numCUs,
                                       cl_uint maxWGSize,
                                       cl_ulong maxAllocSize,
                                       cl_device_type deviceType,
                                       size_t dataTypeSize,
                                       size_t &globalWIs,
                                       size_t &localWIs)
    {
        uint computeDpWgsPerCU; //每个计算单元的工作组数
        if (deviceType & CL_DEVICE_TYPE_CPU)
            computeDpWgsPerCU = 512;
        else {
            // GPU
            computeDpWgsPerCU = 2048;
        }
        globalWIs = numCUs * computeDpWgsPerCU * maxWGSize;
        uint64_t temp = std::min((globalWIs * dataTypeSize), (size_t)maxAllocSize) / dataTypeSize;
        uint64_t n = (temp > UINT64_MAX) ? UINT64_MAX : temp;
        globalWIs = (n / maxWGSize) * maxWGSize;
        localWIs = maxWGSize;
    }
};

template <class T>
class BasicDataPerformance
{
public:
    typedef struct {
        std::string kernelname;             // kernel name
        std::string indexVarInit;     // initialization formula for the index
        std::string opFormula;        // arithmetic formula for the accumulator
        int numStreams;               // number of parallel streams
        int numUnrolls;               // number of times the loop was unrolled
        int numRepeats;               // number of loop iterations (>= 1)
        int flopCount;
    } benchmark_type;

    float perform(cl_device_id id,
                  cl_context context,
                  cl_command_queue queue,
                  const std::string &kernelHeader = "")
    {
        double gflopsall = 0;
        clFlush(queue);
        std::list<benchmark_type> lists = benchmarklist();
        typename  std::list<benchmark_type>::iterator it = lists.begin();
        while (it != lists.end()) {
            cl_event evt;
            cl_int err;
            double runtime = 0;
            size_t globalWIs;
            size_t localWIs;
            size_t workItemSize;
            T *hostMem = nullptr;
            T *hostMemRes = nullptr;
            benchmark_type benmark = *it++;
            CLDeviceInfo deviceinfo(id);

            CLTkit::getNDRangeWorkItemSize(deviceinfo.maxComputeUnit(),
                                           deviceinfo.maxWGSize(),
                                           deviceinfo.maxAllocSize(),
                                           deviceinfo.deviceType(),
                                           sizeof(T),
                                           globalWIs,
                                           localWIs);
            workItemSize = globalWIs;

            hostMem  = new T[workItemSize];
            hostMemRes = new T[workItemSize];

            cl_mem cl_result = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(T) * workItemSize, NULL,
                                              NULL);

            // init host memory
            for (int i = 0; i < workItemSize; ++i) {
                hostMem[i] = (T)(drand48() * 5.0);
            }

            // write buffer from host
            err = clEnqueueWriteBuffer (queue, cl_result, CL_TRUE, 0,
                                        sizeof(T) * workItemSize, hostMem,
                                        0, NULL, &evt);
            clWaitForEvents(1, &evt);
            clReleaseEvent(evt);

            std::string kernelCode = generateKernel(benmark, kernelHeader);
            // Create kernel program object
            const char *progSource[] = {kernelCode.c_str()};
            cl_program program = clCreateProgramWithSource(context, 1, progSource,
                                                           NULL, &err);
            CHECK_CL_ERROR(err);
            // Compile the program
            err = clBuildProgram(program, 1, &id, NULL, NULL, NULL);
            CHECK_CL_ERROR(err);
            // create kernel
            cl_kernel kernel = clCreateKernel(program, benmark.kernelname.c_str(), &err);
            CHECK_CL_ERROR(err);
            clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_result);

            // send to device queue
            size_t global_work_size[1] = {workItemSize};
            size_t local_work_size[1]  = {localWIs};
            cl_event enentPoint, event0;
            err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, NULL, 0,
                                         NULL,
                                         &evt);
            CHECK_CL_ERROR(err);
            clWaitForEvents(1, &evt);
            clFinish(queue);
            // kernel execution time. Milliseconds
            cl_ulong startTime, endTime;
            err = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_SUBMIT,
                                          sizeof(cl_ulong), &startTime, NULL);
            err = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_END,
                                          sizeof(cl_ulong), &endTime, NULL);

            clReleaseEvent(evt);
            runtime =  (endTime -  startTime) / (1000.0 * 1000.0);

            double flopcout = (benmark.numStreams *
                               benmark.numUnrolls *
                               benmark.numRepeats *
                               benmark.flopCount) ;

            double gflops = ((workItemSize * flopcout) /
                             (runtime * 1.0)) * 1e3f;
            gflops /= 1e9f;
            gflopsall += gflops;
            // read from device
            err = clEnqueueReadBuffer(queue, cl_result, CL_TRUE, 0,
                                      workItemSize * sizeof(T), hostMemRes,
                                      0, NULL, &evt);

            clWaitForEvents(1, &evt);
            clFinish(queue);
            clReleaseEvent(evt);
            //release all
            clReleaseKernel(kernel);
            clReleaseProgram(program);
            clReleaseMemObject(cl_result);
            if (hostMem) {
                delete []hostMem;
                hostMem = nullptr;
            }
            if (hostMemRes) {
                delete []hostMemRes;
                hostMemRes = nullptr;
            }
        }
        return gflopsall / lists.size();
    }
private:
    std::list<benchmark_type> benchmarklist()
    {
        std::list<benchmark_type> benchmarklist;
        benchmark_type info[] = {
            {"Add1", "result[gid]", "$+$*$", 1, 128, 20, 2},
            {"Add2", "result[gid]", "$+$*$", 2, 64, 20, 2},
            {"Add4", "result[gid]", "$+$*$", 4, 32, 20, 2},
            {"Add8", "result[gid]", "$+$*$", 8, 16, 20, 2},
            {"Add16", "result[gid]", "$+$*$", 16, 8, 20, 2},
        };
        for (int i = 0; i < sizeof(info) / sizeof(info[0]); ++i) {
            benchmarklist.push_back(info[i]);
        }
        return benchmarklist;
    }

    std::string generateKernel( const  benchmark_type &test,
                                const  std::string &kernelHeader)
    {
#if (__GNUC__ > 5)
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wignored-attributes"
#endif
        std::ostringstream oss;
        const string kName = test.kernelname;
        string tName = "float";
        if (std::is_same<T, cl_float>::value)
            tName    = "float";
        else if (std::is_same<T, cl_double>::value)
            tName    = "double";
        else if (std::is_same<T, cl_half>::value)
            tName    = "half";
        else if (std::is_same<T, cl_int>::value)
            tName    = "int";
        else
            tName    = "float";
#if (__GNUC__ > 5)
#pragma GCC diagnostic pop
#endif
        if (kernelHeader.size() > 0) {
            char bufheader[50];
            memset(&bufheader, 0, 50);
            sprintf(bufheader, "#pragma OPENCL EXTENSION %s: enable", kernelHeader.c_str());
            oss << bufheader << endl;
        } else
            oss << kernelHeader << endl;

        oss << string("__kernel void ") << kName << "(__global " << tName << " *result) \n{\n"
            << "  int gid = get_global_id(0);\n";

        int numVecs[5] = {0, 0, 0, 0, 0};
        int startIdx[5] = {0, 0, 0, 0, 0};
        int i, nStreams = test.numStreams, nVectors = 0;
        string varInit = string(test.indexVarInit);
        int pos = -1;
        while ((pos = varInit.find("T")) != (-1))
            varInit.replace(pos, 1, tName);

        oss << "  " << tName << " " << "temp" << " = " << varInit << ";\n";
        float iniVal = 0.f;
        for (i = 4 ; i >= 0 ; --i) {
            numVecs[i] = nStreams / (1 << i);
            nStreams -= (numVecs[i] * (1 << i));
            if (i == 4) startIdx[i] = 0;
            else startIdx[i] = startIdx[i + 1] + numVecs[i + 1];
            nVectors += numVecs[i];

            for (int vv = startIdx[i] ; vv < startIdx[i] + numVecs[i] ; ++vv) {
                oss << "  " << tName;
                if (i > 0) oss << (1 << i);
                oss << " " << "data" << " = ";
                if (i > 0) oss << "(" << tName << (1 << i) << ")(";
                oss << "temp";
                for (int ss = 1 ; ss < (1 << i) ; ++ss) {
                    oss << "," << "temp" << "+" << ss;
                }
                if (i > 0) oss << ")";
                oss << ";\n";
            }
        }
        // write body
        char buf[50];
        if (test.numRepeats > 1) {
            sprintf (buf, "  for (int i=0; i<%d; ++i) \n  {\n", test.numRepeats);
            oss << buf;
        }
        // write the body of the loop
        if (test.numUnrolls > 1) {
            sprintf (buf, "    for (int j=0; j<%d; ++j) \n    {\n", test.numUnrolls);
            oss << buf;
        }
        for (int ss = 0 ; ss < nVectors ; ++ss) {
            string opCode = string(test.opFormula);
            int pos = -1;
            sprintf (buf, "%s", "data");
            string lVar = string(buf);
            while ((pos = opCode.find("$")) != (-1))
                opCode.replace(pos, 1, lVar);
            oss << "     " << lVar << "=" << opCode << ";";
        }
        oss << "\n";
        if (test.numUnrolls > 1)
            oss << "    }\n";

        if (test.numRepeats > 1)
            oss << "  }\n";
        oss << "  result[gid] = ";
        bool first = true;
        for (i = 4 ; i >= 0 ; --i)
            if (numVecs[i] > 0) { // this is it
                for (int ss = 0 ; ss < (1 << i) ; ++ss) {
                    if (! first) {
                        oss << "+";
                    } else
                        first = false;
                    oss <<  "data";
                    if (i > 0)
                        oss << ".s" << hex << ss << dec;
                }
            }
        oss << ";\n}";

        std::string kernelCode(oss.str());
        return kernelCode;
    }
};


#endif
